{
  "nbformat": 4,
  "nbformat_minor": 0,
  "metadata": {
    "colab": {
      "provenance": [],
      "machine_shape": "hm",
      "gpuType": "T4",
      "authorship_tag": "ABX9TyPltsbvDGcduB31LLcU4Euo",
      "include_colab_link": true
    },
    "kernelspec": {
      "name": "python3",
      "display_name": "Python 3"
    },
    "language_info": {
      "name": "python"
    },
    "accelerator": "GPU"
  },
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "view-in-github",
        "colab_type": "text"
      },
      "source": [
        "<a href=\"https://colab.research.google.com/github/DanielWarfield1/MLWritingAndResearch/blob/main/CUDA_cheet_sheet.ipynb\" target=\"_parent\"><img src=\"https://colab.research.google.com/assets/colab-badge.svg\" alt=\"Open In Colab\"/></a>"
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "# CUDA cheet sheet\n",
        "Runtime must be NVIDIA GPU\n",
        "\n",
        "Based on these amazing tutorials:\n",
        "\n",
        "https://www.youtube.com/watch?v=m0nhePeHwFs&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK&index=2"
      ],
      "metadata": {
        "id": "d3ZIjbcjB-24"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "# Setting up CUDA"
      ],
      "metadata": {
        "id": "Vadfxhq3EmoZ"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "!pip install nvcc4jupyter"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "carVN0NxFtsq",
        "outputId": "95ff191c-28a4-414d-c46c-147a791b983a"
      },
      "execution_count": 1,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Collecting nvcc4jupyter\n",
            "  Downloading nvcc4jupyter-1.2.1-py3-none-any.whl (10 kB)\n",
            "Installing collected packages: nvcc4jupyter\n",
            "Successfully installed nvcc4jupyter-1.2.1\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%load_ext nvcc4jupyter"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "UQQO-a9Pi7rW",
        "outputId": "c3a2c102-e3c6-4e51-aba8-d048a5d67311"
      },
      "execution_count": 2,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Detected platform \"Colab\". Running its setup...\n",
            "Source files will be saved in \"/tmp/tmpd3e4ofbt\".\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "# Hello World in CUDA"
      ],
      "metadata": {
        "id": "LxEEcrNSjH62"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "#include <stdio.h>\n",
        "\n",
        "__global__ void hello(){\n",
        "    printf(\"Hello from block: %u, thread: %u\\n\", blockIdx.x, threadIdx.x);\n",
        "}\n",
        "\n",
        "__host__ int main(){\n",
        "    hello<<<2, 2>>>();\n",
        "    cudaDeviceSynchronize();\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "bbZ3-5YPi9s3",
        "outputId": "9f81b71e-d74c-4474-aba4-375d46dc4f5f"
      },
      "execution_count": 3,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Hello from block: 0, thread: 0\n",
            "Hello from block: 0, thread: 1\n",
            "Hello from block: 1, thread: 0\n",
            "Hello from block: 1, thread: 1\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "#include <stdio.h>\n",
        "\n",
        "__global__ void hello() {\n",
        "    printf(\"Hello from block: (%u, %u, %u), thread: (%u, %u, %u)\\n\",\n",
        "           blockIdx.x, blockIdx.y, blockIdx.z,\n",
        "           threadIdx.x, threadIdx.y, threadIdx.z);\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    // Define the dimensions of the grid and blocks\n",
        "    dim3 gridDim(2, 2, 2);   // 2x2x2 grid of blocks\n",
        "    dim3 blockDim(2, 2, 2);  // 2x2x2 grid of threads per block\n",
        "\n",
        "    // Launch the kernel\n",
        "    hello<<<gridDim, blockDim>>>();\n",
        "\n",
        "    // Wait for GPU to finish before accessing on host\n",
        "    cudaDeviceSynchronize();\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "xwgCAesZJLo-",
        "outputId": "86d0a9db-fd72-44ad-90e9-9b0a4b23060d"
      },
      "execution_count": 4,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Hello from block: (0, 0, 1), thread: (0, 0, 0)\n",
            "Hello from block: (0, 0, 1), thread: (1, 0, 0)\n",
            "Hello from block: (0, 0, 1), thread: (0, 1, 0)\n",
            "Hello from block: (0, 0, 1), thread: (1, 1, 0)\n",
            "Hello from block: (0, 0, 1), thread: (0, 0, 1)\n",
            "Hello from block: (0, 0, 1), thread: (1, 0, 1)\n",
            "Hello from block: (0, 0, 1), thread: (0, 1, 1)\n",
            "Hello from block: (0, 0, 1), thread: (1, 1, 1)\n",
            "Hello from block: (0, 0, 0), thread: (0, 0, 0)\n",
            "Hello from block: (0, 0, 0), thread: (1, 0, 0)\n",
            "Hello from block: (0, 0, 0), thread: (0, 1, 0)\n",
            "Hello from block: (0, 0, 0), thread: (1, 1, 0)\n",
            "Hello from block: (0, 0, 0), thread: (0, 0, 1)\n",
            "Hello from block: (0, 0, 0), thread: (1, 0, 1)\n",
            "Hello from block: (0, 0, 0), thread: (0, 1, 1)\n",
            "Hello from block: (0, 0, 0), thread: (1, 1, 1)\n",
            "Hello from block: (1, 0, 1), thread: (0, 0, 0)\n",
            "Hello from block: (1, 0, 1), thread: (1, 0, 0)\n",
            "Hello from block: (1, 0, 1), thread: (0, 1, 0)\n",
            "Hello from block: (1, 0, 1), thread: (1, 1, 0)\n",
            "Hello from block: (1, 0, 1), thread: (0, 0, 1)\n",
            "Hello from block: (1, 0, 1), thread: (1, 0, 1)\n",
            "Hello from block: (1, 0, 1), thread: (0, 1, 1)\n",
            "Hello from block: (1, 0, 1), thread: (1, 1, 1)\n",
            "Hello from block: (0, 1, 0), thread: (0, 0, 0)\n",
            "Hello from block: (0, 1, 0), thread: (1, 0, 0)\n",
            "Hello from block: (0, 1, 0), thread: (0, 1, 0)\n",
            "Hello from block: (0, 1, 0), thread: (1, 1, 0)\n",
            "Hello from block: (0, 1, 0), thread: (0, 0, 1)\n",
            "Hello from block: (0, 1, 0), thread: (1, 0, 1)\n",
            "Hello from block: (0, 1, 0), thread: (0, 1, 1)\n",
            "Hello from block: (0, 1, 0), thread: (1, 1, 1)\n",
            "Hello from block: (1, 1, 1), thread: (0, 0, 0)\n",
            "Hello from block: (1, 1, 1), thread: (1, 0, 0)\n",
            "Hello from block: (1, 1, 1), thread: (0, 1, 0)\n",
            "Hello from block: (1, 1, 1), thread: (1, 1, 0)\n",
            "Hello from block: (1, 1, 1), thread: (0, 0, 1)\n",
            "Hello from block: (1, 1, 1), thread: (1, 0, 1)\n",
            "Hello from block: (1, 1, 1), thread: (0, 1, 1)\n",
            "Hello from block: (1, 1, 1), thread: (1, 1, 1)\n",
            "Hello from block: (1, 1, 0), thread: (0, 0, 0)\n",
            "Hello from block: (1, 1, 0), thread: (1, 0, 0)\n",
            "Hello from block: (1, 1, 0), thread: (0, 1, 0)\n",
            "Hello from block: (1, 1, 0), thread: (1, 1, 0)\n",
            "Hello from block: (1, 1, 0), thread: (0, 0, 1)\n",
            "Hello from block: (1, 1, 0), thread: (1, 0, 1)\n",
            "Hello from block: (1, 1, 0), thread: (0, 1, 1)\n",
            "Hello from block: (1, 1, 0), thread: (1, 1, 1)\n",
            "Hello from block: (1, 0, 0), thread: (0, 0, 0)\n",
            "Hello from block: (1, 0, 0), thread: (1, 0, 0)\n",
            "Hello from block: (1, 0, 0), thread: (0, 1, 0)\n",
            "Hello from block: (1, 0, 0), thread: (1, 1, 0)\n",
            "Hello from block: (1, 0, 0), thread: (0, 0, 1)\n",
            "Hello from block: (1, 0, 0), thread: (1, 0, 1)\n",
            "Hello from block: (1, 0, 0), thread: (0, 1, 1)\n",
            "Hello from block: (1, 0, 0), thread: (1, 1, 1)\n",
            "Hello from block: (0, 1, 1), thread: (0, 0, 0)\n",
            "Hello from block: (0, 1, 1), thread: (1, 0, 0)\n",
            "Hello from block: (0, 1, 1), thread: (0, 1, 0)\n",
            "Hello from block: (0, 1, 1), thread: (1, 1, 0)\n",
            "Hello from block: (0, 1, 1), thread: (0, 0, 1)\n",
            "Hello from block: (0, 1, 1), thread: (1, 0, 1)\n",
            "Hello from block: (0, 1, 1), thread: (0, 1, 1)\n",
            "Hello from block: (0, 1, 1), thread: (1, 1, 1)\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "# First Kernel\n",
        "https://www.youtube.com/watch?v=rbHURDt2dcQ&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK&index=2\n",
        "\n",
        "---\n",
        "\n",
        "a kernel is a function that does something on a GPU.\n",
        "\n",
        "---\n",
        "\n",
        "`__global__` is a kernel. It's called by the host (cpu) but it runs on the device (gpu)\n",
        "\n",
        "---\n",
        "\n",
        "`cudaMalloc` is a function that reserves global memory on the GPU. It takes two arguments:\n",
        "- a pointer on the device (GPU)\n",
        "- some ammount of space on the GPU in bytes.\n",
        "\n",
        "---\n",
        "\n",
        "`cudaFree` frees up memory on the device. Accepts a pointer\n",
        "\n",
        "---\n",
        "\n",
        "`cudaMemcpy` copies data from the host (computer RAM) to the device (GPU memory) and vice versa. Accepts four arguments\n",
        "- a pointer for the destination\n",
        "- a pointer for the source\n",
        "- a size in bytes\n",
        "- a direction (for instance, from the device to the host, or from the host to device).\n",
        "\n",
        "Keep in mind, this is highly efficient C++. That means we don't get a lot of abstractions to help us out.\n",
        "\n",
        "---\n",
        "\n",
        "First try, getting GPU to add two numbers together:\n",
        "- Create two ints\n",
        "- allocate mem on device\n",
        "- copy values\n",
        "- call a kernel to add them\n",
        "- copy back\n",
        "- print out results\n",
        "- free device memory"
      ],
      "metadata": {
        "id": "b3rTmXV2jmSu"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "#include <iostream>\n",
        "#include <cuda.h>\n",
        "\n",
        "using namespace std;\n",
        "\n",
        "// Defining the kernel\n",
        "__global__ void addIntsCUDA(int *a, int *b) {\n",
        "    a[0] += b[0];\n",
        "}\n",
        "\n",
        "// Running main on host, which triggers the kernel\n",
        "int main() {\n",
        "    // Host values\n",
        "    int a = 1, b = 2;\n",
        "\n",
        "    //printing expression\n",
        "    cout << a << \" + \" << b <<\" = \";\n",
        "\n",
        "    // Device pointers (GPU)\n",
        "    int *d_a, *d_b;\n",
        "\n",
        "    // Allocating memory on the device (GPU)\n",
        "    cudaMalloc(&d_a, sizeof(int));\n",
        "    cudaMalloc(&d_b, sizeof(int));\n",
        "\n",
        "    // Copying values from the host (CPU RAM) to the device (GPU)\n",
        "    cudaMemcpy(d_a, &a, sizeof(int), cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_b, &b, sizeof(int), cudaMemcpyHostToDevice);\n",
        "\n",
        "    // Calling the kernel to add the two values at the two pointer locations.\n",
        "    addIntsCUDA<<<1, 1>>>(d_a, d_b);\n",
        "\n",
        "    // The addition function overwrites the a pointer with the sum. Thus\n",
        "    // this copies the result.\n",
        "    cudaMemcpy(&a, d_a, sizeof(int), cudaMemcpyDeviceToHost);\n",
        "\n",
        "    //printing result\n",
        "    cout << a << endl;\n",
        "\n",
        "    //freeing memory.\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "wYXPpeWDjEQS",
        "outputId": "ed62616c-bdd3-43e1-b257-b501249e2221"
      },
      "execution_count": 5,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "1 + 2 = 3\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "# Multi Threaded Addition\n",
        "https://www.youtube.com/watch?v=kzXjRFL-gjo&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK&index=4\n",
        "\n",
        "---\n",
        "\n",
        "`thead` is a single execution unit that runs a kernel.\n",
        "\n",
        "---\n",
        "\n",
        "`thread block` all the threads in a block can communicate with one another.\n",
        "\n",
        "---\n",
        "\n",
        "`grid` is a bunch of thread blocks.\n",
        "\n",
        "---\n",
        "\n",
        "`dim3` a 3d structure in CUDA\n",
        "not super clear exactly what this is in the video. Will be elaborated more\n",
        "\n",
        "---\n",
        "\n",
        "`launch configuration` something like `<<<100,256>>>` would launch 100 blocks of 256 threads. One invokes a kernel with a launch configuration. Like `myKernel<<<100, 256>>>(arguments)`. It seems like you can call 1024 threads per block, and you can call a lot of blocks (like 2^31-1)\n",
        "\n",
        "There's a limit on the number of threads per block, but an almost infinite number of blocks. A big GPU might run multiple blocks at once, a weaker GPU might only run one or two blocks at once. This makes high powered GPUs run similarly to low powered GPUs, just faster.\n",
        "\n",
        "---\n",
        "\n",
        "each thread is an individula with some unique identifiers that allow the threads to do certain things\n",
        "\n",
        "- threadIdx\n",
        "- blockIdx\n",
        "- blockDim\n",
        "- gridDim\n",
        "\n",
        "these allow a thread to get an idea of which thread it is within a block, and which block it exists in within a set of blocks.\n",
        "\n",
        "---\n",
        "\n",
        "We can use this to parallelize a whole bunch of addition.\n",
        "\n",
        "\n"
      ],
      "metadata": {
        "id": "BJit2yjbtwuI"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "#include <iostream>\n",
        "#include <cuda.h>\n",
        "#include <stdlib.h>\n",
        "#include <ctime>\n",
        "#include <chrono>\n",
        "\n",
        "using namespace std::chrono;\n",
        "using namespace std;\n",
        "\n",
        "__global__ void addInts(int *a, int *b, int count){\n",
        "    int id = blockIdx.x* blockDim.x + threadIdx.x;\n",
        "    if(id < count){\n",
        "        a[id] += b[id];\n",
        "    }\n",
        "}\n",
        "\n",
        "int main(){\n",
        "    srand(time(NULL));\n",
        "    int count = 500000000;\n",
        "    int *h_a = new int[count];\n",
        "    int *h_b = new int[count];\n",
        "\n",
        "    for (int i = 0; i < count; i++){\n",
        "        h_a[i] = rand() % 1000;\n",
        "        h_b[i] = rand() % 1000;\n",
        "    }\n",
        "\n",
        "    auto start = high_resolution_clock::now();\n",
        "\n",
        "    cout<<\"First 5 Calculations out of \"<<count<<\" : \"<<endl;\n",
        "    for(int i = 0; i < 5; i++){\n",
        "        cout<<h_a[i]<<\" + \"<<h_b[i]<<endl;\n",
        "    }\n",
        "\n",
        "    int *d_a, *d_b;\n",
        "\n",
        "    if(cudaMalloc(&d_a, sizeof(int) * count) != cudaSuccess){\n",
        "        cout<<\"error allocating a\";\n",
        "        return 0;\n",
        "    }\n",
        "\n",
        "    if(cudaMalloc(&d_b, sizeof(int) * count) != cudaSuccess){\n",
        "        cout<<\"error allocating b\";\n",
        "        cudaFree(d_a);\n",
        "        delete[] h_a;\n",
        "        delete[] h_b;\n",
        "        return 0;\n",
        "    }\n",
        "\n",
        "    if(cudaMemcpy(d_a, h_a, sizeof(int) * count, cudaMemcpyHostToDevice) != cudaSuccess){\n",
        "        cout<<\"could not copy\"<<endl;\n",
        "        cudaFree(d_a);\n",
        "        cudaFree(d_b);\n",
        "        delete[] h_a;\n",
        "        delete[] h_b;\n",
        "        return 0;\n",
        "    }\n",
        "\n",
        "    if(cudaMemcpy(d_b, h_b, sizeof(int) * count, cudaMemcpyHostToDevice) != cudaSuccess){\n",
        "        cout<<\"could not copy\"<<endl;\n",
        "        cudaFree(d_a);\n",
        "        cudaFree(d_b);\n",
        "        delete[] h_a;\n",
        "        delete[] h_b;\n",
        "        return 0;\n",
        "    }\n",
        "\n",
        "    auto calcStart = high_resolution_clock::now();\n",
        "    addInts<<<count/1024+1, 1024>>>(d_a, d_b, count);\n",
        "    auto calcStop = high_resolution_clock::now();\n",
        "\n",
        "    if(cudaMemcpy(h_a, d_a, sizeof(int)*count, cudaMemcpyDeviceToHost) !=cudaSuccess){\n",
        "        cudaFree(d_a);\n",
        "        cudaFree(d_b);\n",
        "        delete[] h_a;\n",
        "        delete[] h_b;\n",
        "        cout<<\"error copying from device\"<<endl;\n",
        "        return 0;\n",
        "    }\n",
        "\n",
        "    auto stop = high_resolution_clock::now();\n",
        "    auto duration = duration_cast<microseconds>(stop - start);\n",
        "\n",
        "    cout<<endl<<\"Results from GPU:\"<<endl;\n",
        "    for(int i = 0; i < 5; i++){\n",
        "        cout<<h_a[i]<<endl;\n",
        "    }\n",
        "\n",
        "    cout << endl << \"time to do everything: \" << duration.count() << \" microseconds\" << endl;\n",
        "    auto calcDuration = duration_cast<microseconds>(calcStop - calcStart);\n",
        "    cout << endl << \"time to do only calculation: \" << calcDuration.count() << \" microseconds\" << endl;\n",
        "\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "    delete[] h_a;\n",
        "    delete[] h_b;\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "KjGnnF_AqKod",
        "outputId": "35ae2255-3863-4e12-a20e-324b9cf86793"
      },
      "execution_count": 6,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "First 5 Calculations out of 500000000 : \n",
            "841 + 299\n",
            "944 + 447\n",
            "395 + 755\n",
            "952 + 235\n",
            "367 + 421\n",
            "\n",
            "Results from GPU:\n",
            "1140\n",
            "1391\n",
            "1150\n",
            "1187\n",
            "788\n",
            "\n",
            "time to do everything: 1427178 microseconds\n",
            "\n",
            "time to do only calculation: 39719 microseconds\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## Comparing to CPU"
      ],
      "metadata": {
        "id": "od9NtDlO618X"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "#include <iostream>\n",
        "#include <thread>\n",
        "#include <vector>\n",
        "#include <stdlib.h>\n",
        "#include <ctime>\n",
        "#include <chrono>\n",
        "\n",
        "using namespace std::chrono;\n",
        "using namespace std;\n",
        "\n",
        "void addInts(int *a, int *b, int start, int end) {\n",
        "    for (int i = start; i < end; ++i) {\n",
        "        a[i] += b[i];\n",
        "    }\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    srand(time(NULL));\n",
        "    int count = 500000000;\n",
        "    int *h_a = new int[count];\n",
        "    int *h_b = new int[count];\n",
        "\n",
        "    for (int i = 0; i < count; i++) {\n",
        "        h_a[i] = rand() % 1000;\n",
        "        h_b[i] = rand() % 1000;\n",
        "    }\n",
        "\n",
        "    auto start = high_resolution_clock::now();\n",
        "\n",
        "    cout << \"First 5 Calculations out of \" << count << \" : \" << endl;\n",
        "    for (int i = 0; i < 5; i++) {\n",
        "        cout << h_a[i] << \" + \" << h_b[i] << endl;\n",
        "    }\n",
        "\n",
        "    int numThreads = 8;\n",
        "    vector<thread> threads;\n",
        "    int chunkSize = count / numThreads;\n",
        "\n",
        "    for (int i = 0; i < numThreads; ++i) {\n",
        "        int startIdx = i * chunkSize;\n",
        "        int endIdx = (i == numThreads - 1) ? count : startIdx + chunkSize;\n",
        "        threads.push_back(thread(addInts, h_a, h_b, startIdx, endIdx));\n",
        "    }\n",
        "\n",
        "    for (auto &t : threads) {\n",
        "        t.join();\n",
        "    }\n",
        "\n",
        "    auto stop = high_resolution_clock::now();\n",
        "    auto duration = duration_cast<microseconds>(stop - start);\n",
        "\n",
        "    cout << endl << \"Results from CPU:\" << endl;\n",
        "    for (int i = 0; i < 5; i++) {\n",
        "        cout << h_a[i] << endl;\n",
        "    }\n",
        "\n",
        "    cout << endl << \"time to do all calculations: \" << duration.count() << \" microseconds\" << endl;\n",
        "\n",
        "    delete[] h_a;\n",
        "    delete[] h_b;\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "5o97-UI41Uwm",
        "outputId": "683ca7fd-2c1a-448d-d40b-d21c43e6bde4"
      },
      "execution_count": 7,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "First 5 Calculations out of 500000000 : \n",
            "638 + 932\n",
            "663 + 716\n",
            "385 + 660\n",
            "874 + 113\n",
            "124 + 716\n",
            "\n",
            "Results from CPU:\n",
            "1570\n",
            "1379\n",
            "1045\n",
            "987\n",
            "840\n",
            "\n",
            "time to do all calculations: 310649 microseconds\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "example of how ludicrously long it takes a GPU to get data vs how fast the calculations get done."
      ],
      "metadata": {
        "id": "VVxA5Shv8okT"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "## Memory in CUDA\n",
        "https://www.youtube.com/watch?v=RY2_8wB2QY4&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK&index=5\n",
        "\n",
        "---\n",
        "\n",
        "**note this is an old video, so some of this might be pretty different**\n",
        "\n",
        "Host can read or write to:\n",
        "- Global\n",
        "- Constant\n",
        "- Texture\n",
        "\n",
        "All blocks can read or write from Global, constant (?) or texture (?)\n",
        "\n",
        "Blocks have shared memory that's accessible to each thread, each thread has local memory and registers.\n",
        "\n",
        "Global memory:\n",
        "- Global is big and slow\n",
        "- cudaMalloc, cudaFree, cudaMemcpy, cudaMemset all deal with global memory\n",
        "- This is the main memory of the graphics card.\n",
        "- GPU and CPU can address every byte\n",
        "- persistant across kernel calls\n",
        "\n",
        "Local Memory:\n",
        "- also slow\n",
        "- used when registers can't be used, or when we run out\n",
        "- used for arrays and big objects\n",
        "- Cached (L1 and L2 cache), so the compiler tries it's best to manage this memory to preserve performance.\n",
        "\n",
        "Caches:\n",
        "- L1 cache is really fast (basically a partition of shared memory)\n",
        "- L2 is slower, all shared memory go through L2\n",
        "- you can turn caching on and off\n",
        "\n",
        "Constant Memory:\n",
        "- part of GPU main memory\n",
        "- very close to threads, so fast\n",
        "- CPU can write, GPU can only read\n",
        "- very fast\n",
        "- not a lot of it (in this old 10 year old tutorial)\n",
        "\n",
        "Texture Memory:\n",
        "- Read only by GPU, set up by CPU\n",
        "- designed to do tecture fetching\n",
        "- interpolates pixels\n",
        "- cached in a way that there are quirks\n",
        "\n",
        "Shared memory\n",
        "- very fast\n",
        "- try not to serialize, want all threads to access at the same time.\n",
        "- an L1 cache that we're in control of\n",
        "\n",
        "Registers\n",
        "- basically a variable. int i=5, for instance. Unless we run out, then local memory will be used\n",
        "- scope is per thread\n",
        "- Streaming multiprocessors (sms) can execute more than one thread per time, if resources permit (?) so being light on register use is useful\n",
        "\n",
        "breakdown image at 14:36"
      ],
      "metadata": {
        "id": "E6vKttNq9PHW"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "\n",
        "#include <iostream>\n",
        "#include <cuda_runtime.h>\n",
        "\n",
        "int main() {\n",
        "    int deviceCount;\n",
        "    cudaGetDeviceCount(&deviceCount);\n",
        "\n",
        "    std::cout << \"Number of CUDA devices: \" << deviceCount << std::endl;\n",
        "\n",
        "    for (int device = 0; device < deviceCount; ++device) {\n",
        "        cudaDeviceProp deviceProp;\n",
        "        cudaGetDeviceProperties(&deviceProp, device);\n",
        "\n",
        "        std::cout << \"\\nDevice \" << device << \": \\\"\" << deviceProp.name << \"\\\"\" << std::endl;\n",
        "        std::cout << \"  Total amount of global memory: \" << deviceProp.totalGlobalMem << \" bytes\" << std::endl;\n",
        "        std::cout << \"  Total number of multiprocessors: \" << deviceProp.multiProcessorCount << std::endl;\n",
        "        std::cout << \"  Total amount of constant memory: \" << deviceProp.totalConstMem << \" bytes\" << std::endl;\n",
        "        std::cout << \"  Total amount of shared memory per block: \" << deviceProp.sharedMemPerBlock << \" bytes\" << std::endl;\n",
        "        std::cout << \"  Total number of registers available per block: \" << deviceProp.regsPerBlock << std::endl;\n",
        "        std::cout << \"  Warp size: \" << deviceProp.warpSize << std::endl;\n",
        "        std::cout << \"  Maximum number of threads per block: \" << deviceProp.maxThreadsPerBlock << std::endl;\n",
        "        std::cout << \"  Maximum sizes of each dimension of a block: \" << deviceProp.maxThreadsDim[0] << \" x \" << deviceProp.maxThreadsDim[1] << \" x \" << deviceProp.maxThreadsDim[2] << std::endl;\n",
        "        std::cout << \"  Maximum sizes of each dimension of a grid: \" << deviceProp.maxGridSize[0] << \" x \" << deviceProp.maxGridSize[1] << \" x \" << deviceProp.maxGridSize[2] << std::endl;\n",
        "        std::cout << \"  Clock rate: \" << deviceProp.clockRate << \" kilohertz\" << std::endl;\n",
        "        std::cout << \"  Memory clock rate: \" << deviceProp.memoryClockRate << \" kilohertz\" << std::endl;\n",
        "        std::cout << \"  Memory bus width: \" << deviceProp.memoryBusWidth << \" bits\" << std::endl;\n",
        "    }\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "sg0caePu67BK",
        "outputId": "2cd172c9-1060-40d7-ec50-dd915dc425bb"
      },
      "execution_count": 8,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Number of CUDA devices: 1\n",
            "\n",
            "Device 0: \"Tesla T4\"\n",
            "  Total amount of global memory: 15835660288 bytes\n",
            "  Total number of multiprocessors: 40\n",
            "  Total amount of constant memory: 65536 bytes\n",
            "  Total amount of shared memory per block: 49152 bytes\n",
            "  Total number of registers available per block: 65536\n",
            "  Warp size: 32\n",
            "  Maximum number of threads per block: 1024\n",
            "  Maximum sizes of each dimension of a block: 1024 x 1024 x 64\n",
            "  Maximum sizes of each dimension of a grid: 2147483647 x 65535 x 65535\n",
            "  Clock rate: 1590000 kilohertz\n",
            "  Memory clock rate: 5001000 kilohertz\n",
            "  Memory bus width: 256 bits\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## Playing with the Nvidia Profiler\n",
        "- writes some code to a file\n",
        "- compiles it\n",
        "- runs the profiler"
      ],
      "metadata": {
        "id": "F4cGmzSERFX4"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile addInts.cu\n",
        "#include <iostream>\n",
        "#include <cuda.h>\n",
        "#include <stdlib.h>\n",
        "#include <ctime>\n",
        "\n",
        "using namespace std;\n",
        "\n",
        "__global__ void addInts(int *a, int *b, int count){\n",
        "    int id = blockIdx.x* blockDim.x + threadIdx.x;\n",
        "    if(id < count){\n",
        "        a[id] += b[id];\n",
        "    }\n",
        "}\n",
        "\n",
        "int main(){\n",
        "    srand(time(NULL));\n",
        "    int count = 10000;\n",
        "    int *h_a = new int[count];\n",
        "    int *h_b = new int[count];\n",
        "\n",
        "    for (int i = 0; i < count; i++){\n",
        "        h_a[i] = rand() % 1000;\n",
        "        h_b[i] = rand() % 1000;\n",
        "    }\n",
        "\n",
        "    int *d_a, *d_b;\n",
        "\n",
        "    cudaMalloc(&d_a, sizeof(int) * count);\n",
        "    cudaMalloc(&d_b, sizeof(int) * count);\n",
        "\n",
        "    cudaMemcpy(d_a, h_a, sizeof(int) * count, cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_b, h_b, sizeof(int) * count, cudaMemcpyHostToDevice);\n",
        "\n",
        "    addInts<<<count/256+1, 256>>>(d_a, d_b, count);\n",
        "\n",
        "    cudaMemcpy(h_a, d_a, sizeof(int) * count, cudaMemcpyDeviceToHost);\n",
        "\n",
        "    cudaFree(d_a);\n",
        "    cudaFree(d_b);\n",
        "    delete[] h_a;\n",
        "    delete[] h_b;\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "x6D7_iTtOOYQ",
        "outputId": "fb296415-5d54-4a9e-dba2-99c946b7934a"
      },
      "execution_count": 9,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Writing addInts.cu\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc addInts.cu\n",
        "!./a.out"
      ],
      "metadata": {
        "id": "qmwsJ8sFROga"
      },
      "execution_count": 10,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc addInts.cu -o addInts\n",
        "!nvprof ./addInts"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "licAxCkaRRQY",
        "outputId": "15991a26-dff2-4437-a777-08cb4a71fa1d"
      },
      "execution_count": 11,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "==2277== NVPROF is profiling process 2277, command: ./addInts\n",
            "==2277== Profiling application: ./addInts\n",
            "==2277== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   58.91%  12.160us         2  6.0800us  6.0800us  6.0800us  [CUDA memcpy HtoD]\n",
            "                   24.34%  5.0240us         1  5.0240us  5.0240us  5.0240us  [CUDA memcpy DtoH]\n",
            "                   16.74%  3.4560us         1  3.4560us  3.4560us  3.4560us  addInts(int*, int*, int)\n",
            "      API calls:   99.34%  84.915ms         2  42.458ms  4.4530us  84.911ms  cudaMalloc\n",
            "                    0.24%  208.04us         1  208.04us  208.04us  208.04us  cudaLaunchKernel\n",
            "                    0.16%  137.74us       114  1.2080us     144ns  53.741us  cuDeviceGetAttribute\n",
            "                    0.12%  101.00us         2  50.500us  7.0930us  93.908us  cudaFree\n",
            "                    0.11%  89.781us         3  29.927us  22.937us  33.805us  cudaMemcpy\n",
            "                    0.02%  13.175us         1  13.175us  13.175us  13.175us  cuDeviceGetName\n",
            "                    0.01%  5.4920us         1  5.4920us  5.4920us  5.4920us  cuDeviceGetPCIBusId\n",
            "                    0.00%  3.9750us         1  3.9750us  3.9750us  3.9750us  cuDeviceTotalMem\n",
            "                    0.00%  1.7200us         3     573ns     197ns  1.2890us  cuDeviceGetCount\n",
            "                    0.00%     812ns         2     406ns     150ns     662ns  cuDeviceGet\n",
            "                    0.00%     267ns         1     267ns     267ns     267ns  cuModuleGetLoadingMode\n",
            "                    0.00%     217ns         1     217ns     217ns     217ns  cuDeviceGetUuid\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "# Implementing Embarassingly Parallel algorithms\n",
        "\n",
        "yes that's a term, that's when threads don't need to interact at all. (very easy to make parallel). also called pleasingly parallel. We like embarassingly parallel.\n",
        "\n",
        "https://www.youtube.com/watch?v=0ILeCeaor0A&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK&index=6\n",
        "\n",
        "Realistically you have some C++ code, then you want to parallelize it in CUDA. This takes a brute force of nearest neighbor and parallelizes it.\n",
        "\n",
        "\n"
      ],
      "metadata": {
        "id": "kjvW-G2_VLpd"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%cuda\n",
        "#include <iostream>\n",
        "#include <ctime>\n",
        "#include <cuda.h>\n",
        "#include <cuda_runtime.h>\n",
        "#include <device_launch_parameters.h>\n",
        "\n",
        "using namespace std;\n",
        "\n",
        "//brute force implementation on the CPU\n",
        "void findClosestCPU(float3* points, int* indices, int count) {\n",
        "    // Base case, if there's 1 point don't do anything\n",
        "    if(count <=1) return;\n",
        "    // Loop through every point\n",
        "    for (int curPoint = 0; curPoint < count; curPoint++) {\n",
        "        // This variable is nearest so far, set it to float. max\n",
        "        float distToClosest = 3.4028238f ;\n",
        "        // See how far itais from every other point\n",
        "        for (int i = 0; i < count; i++) {\n",
        "            // Don't check distance to itself\n",
        "            if(i == curPoint) continue;\n",
        "            float dist_sqr = (points[curPoint].x - points[i].x) *\n",
        "                (points[curPoint].x - points[i].x) +\n",
        "                (points[curPoint].y - points[i].y) *\n",
        "                (points[curPoint].y - points[i].y) +\n",
        "                (points[curPoint].z - points[i].z) *\n",
        "                (points[curPoint].z - points[i].z);\n",
        "            if(dist_sqr < distToClosest) {\n",
        "                distToClosest = dist_sqr;\n",
        "                indices[curPoint] = i;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "int main(){\n",
        "\n",
        "    //defining parameters\n",
        "    const int count = 10000;\n",
        "    int* indexOfClosest = new int[count];\n",
        "    float3* points = new float3[count];\n",
        "\n",
        "    //defining random points\n",
        "    for (int i = 0; i < count; i++){\n",
        "        points[i].x = (float)(((rand()%10000))-5000);\n",
        "        points[i].y = (float)(((rand()%10000))-5000);\n",
        "        points[i].z = (float)(((rand()%10000))-5000);\n",
        "    }\n",
        "\n",
        "    long fastest = 1000000000;\n",
        "\n",
        "    cout << \"running brute force nearest neighbor on the CPU...\"<<endl;\n",
        "    for (int i = 0; i <= 10; i++){\n",
        "        long start = clock();\n",
        "        findClosestCPU(points, indexOfClosest, count);\n",
        "        double duration = ( clock() - start ) / (double) CLOCKS_PER_SEC;\n",
        "        cout << \"test \" << i << \" took \" << duration << \" seconds\" <<endl;\n",
        "    }\n",
        "\n",
        "    return 0;\n",
        "}\n"
      ],
      "metadata": {
        "id": "yWBXMTJiRTsy",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "0b0fec97-4a08-4e96-8d10-4061f8108a48"
      },
      "execution_count": 12,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "running brute force nearest neighbor on the CPU...\n",
            "test 0 took 1.17471 seconds\n",
            "test 1 took 1.18766 seconds\n",
            "test 2 took 1.18264 seconds\n",
            "test 3 took 1.20396 seconds\n",
            "test 4 took 1.18683 seconds\n",
            "test 5 took 1.20544 seconds\n",
            "test 6 took 1.19764 seconds\n",
            "test 7 took 1.17379 seconds\n",
            "test 8 took 1.19515 seconds\n",
            "test 9 took 1.16731 seconds\n",
            "test 10 took 1.17981 seconds\n",
            "\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "# Implementing on GPU\n",
        "https://www.youtube.com/watch?v=UJ3kNwoxtuk&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK&index=7\n"
      ],
      "metadata": {
        "id": "hRltZ9bRAHdr"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile findClosestGPU.cu\n",
        "#include <iostream>\n",
        "#include <ctime>\n",
        "#include <cuda.h>\n",
        "#include <cuda_runtime.h>\n",
        "#include <device_launch_parameters.h>\n",
        "\n",
        "using namespace std;\n",
        "\n",
        "//this implementation is technically faster, but uses shared memory in an undefined\n",
        "//way, and is thus fairly advanced.\n",
        "\n",
        "//brute force implementation on the CPU\n",
        "__global__ void findClosestGPU(float3* points, int* indices, int count) {\n",
        "    if(count <= 1 ) return;\n",
        "    int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    if (idx < count){\n",
        "        float3 thisPoint = points[idx];\n",
        "        float smallestSoFar = 3.40282e38f;\n",
        "\n",
        "        for (int i = 0; i < count; i++){\n",
        "            if(i == idx) continue;\n",
        "\n",
        "            float dist_sqr = (thisPoint.x - points [i].x) *\n",
        "                (thisPoint.x - points[i].x) +\n",
        "                (thisPoint.y - points[i].y) *\n",
        "                (thisPoint.y - points[i].y) +\n",
        "                (thisPoint.z - points[i].z) *\n",
        "                (thisPoint.z - points[i].z);\n",
        "            if(dist_sqr < smallestSoFar) {\n",
        "                smallestSoFar = dist_sqr;\n",
        "                indices[idx] = i;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "int main(){\n",
        "\n",
        "    //defining parameters\n",
        "    const int count = 10000;\n",
        "    int* d_indexOfClosest = new int[count];\n",
        "    float3* d_points = new float3[count];\n",
        "\n",
        "    //defining random points\n",
        "    for (int i = 0; i < count; i++){\n",
        "        d_points[i].x = (float)(((rand()%10000))-5000);\n",
        "        d_points[i].y = (float)(((rand()%10000))-5000);\n",
        "        d_points[i].z = (float)(((rand()%10000))-5000);\n",
        "    }\n",
        "\n",
        "    int threads_per_block = 64;\n",
        "\n",
        "    cout << \"running brute force nearest neighbor on the CPU...\"<<endl;\n",
        "    for (int i = 1; i <= 10; i++){\n",
        "        long start = clock();\n",
        "\n",
        "        findClosestGPU<<<(count/threads_per_block)+1, threads_per_block>>>(d_points, d_indexOfClosest, count);\n",
        "        cudaMemcpy(d_indexOfClosest, d_indexOfClosest, sizeof(int) * count, cudaMemcpyDeviceToHost);\n",
        "\n",
        "        double duration = ( clock() - start ) / (double) CLOCKS_PER_SEC;\n",
        "        cout << \"test \" << i << \" took \" << duration << \" seconds\" <<endl;\n",
        "    }\n",
        "\n",
        "    return 0;\n",
        "}"
      ],
      "metadata": {
        "id": "NRTExppOdgif",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "57de8720-cea2-4f4c-b9d5-4616f0427484"
      },
      "execution_count": 13,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Writing findClosestGPU.cu\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc findClosestGPU.cu -o findClosestGPU.out\n",
        "!./findClosestGPU.out"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "G_PkqhaiM_cG",
        "outputId": "2c259b65-14c6-4cf9-eda7-c81e09bf8958"
      },
      "execution_count": 14,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "running brute force nearest neighbor on the CPU...\n",
            "test 1 took 0.147381 seconds\n",
            "test 2 took 8e-06 seconds\n",
            "test 3 took 4e-06 seconds\n",
            "test 4 took 4e-06 seconds\n",
            "test 5 took 3e-06 seconds\n",
            "test 6 took 3e-06 seconds\n",
            "test 7 took 4e-06 seconds\n",
            "test 8 took 4e-06 seconds\n",
            "test 9 took 4e-06 seconds\n",
            "test 10 took 3e-06 seconds\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvprof ./findClosestGPU.out"
      ],
      "metadata": {
        "id": "aAGT63l0NB5G",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "ca5c6da8-9fd0-4530-ea2c-5e080e2b275e"
      },
      "execution_count": 15,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "running brute force nearest neighbor on the CPU...\n",
            "==2430== NVPROF is profiling process 2430, command: ./findClosestGPU.out\n",
            "test 1 took 0.102362 seconds\n",
            "test 2 took 1.2e-05 seconds\n",
            "test 3 took 8e-06 seconds\n",
            "test 4 took 9e-06 seconds\n",
            "test 5 took 7e-06 seconds\n",
            "test 6 took 7e-06 seconds\n",
            "test 7 took 6e-06 seconds\n",
            "test 8 took 7e-06 seconds\n",
            "test 9 took 7e-06 seconds\n",
            "test 10 took 8e-06 seconds\n",
            "==2430== Profiling application: ./findClosestGPU.out\n",
            "==2430== Warning: 10 records have invalid timestamps due to insufficient device buffer space. You can configure the buffer space using the option --device-buffer-size.\n",
            "==2430== Profiling result:\n",
            "No kernels were profiled.\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            "      API calls:   99.75%  71.863ms        10  7.1863ms  4.5900us  71.812ms  cudaLaunchKernel\n",
            "                    0.19%  137.44us       114  1.2050us     142ns  58.460us  cuDeviceGetAttribute\n",
            "                    0.02%  14.921us        10  1.4920us     805ns  6.4840us  cudaMemcpy\n",
            "                    0.02%  11.798us         1  11.798us  11.798us  11.798us  cuDeviceGetName\n",
            "                    0.01%  5.3740us         1  5.3740us  5.3740us  5.3740us  cuDeviceGetPCIBusId\n",
            "                    0.01%  3.9970us         1  3.9970us  3.9970us  3.9970us  cuDeviceTotalMem\n",
            "                    0.00%  1.7250us         3     575ns     195ns  1.2710us  cuDeviceGetCount\n",
            "                    0.00%     768ns         2     384ns     147ns     621ns  cuDeviceGet\n",
            "                    0.00%     362ns         1     362ns     362ns     362ns  cuModuleGetLoadingMode\n",
            "                    0.00%     225ns         1     225ns     225ns     225ns  cuDeviceGetUuid\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile findClosestGPU.cu\n",
        "#include <iostream>\n",
        "#include <ctime>\n",
        "#include <cuda.h>\n",
        "#include <cuda_runtime.h>\n",
        "#include <device_launch_parameters.h>\n",
        "\n",
        "using namespace std;\n",
        "\n",
        "// Brute force implementation, parallelized on the GPU\n",
        "__global__ void findClosestGPU(float3* points, int* indices, int count) {\n",
        "    if (count <= 1) return;\n",
        "    int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "    if (idx < count) {\n",
        "        float3 thisPoint = points[idx];\n",
        "        float smallestSoFar = 3.40282e38f;\n",
        "\n",
        "        for (int i = 0; i < count; i++) {\n",
        "            if (i == idx) continue;\n",
        "\n",
        "            float dist_sqr = (thisPoint.x - points[i].x) *\n",
        "                             (thisPoint.x - points[i].x) +\n",
        "                             (thisPoint.y - points[i].y) *\n",
        "                             (thisPoint.y - points[i].y) +\n",
        "                             (thisPoint.z - points[i].z) *\n",
        "                             (thisPoint.z - points[i].z);\n",
        "            if (dist_sqr < smallestSoFar) {\n",
        "                smallestSoFar = dist_sqr;\n",
        "                indices[idx] = i;\n",
        "            }\n",
        "        }\n",
        "    }\n",
        "}\n",
        "\n",
        "int main() {\n",
        "    // Defining parameters\n",
        "    const int count = 10000;\n",
        "    int* h_indexOfClosest = new int[count];\n",
        "    float3* h_points = new float3[count];\n",
        "\n",
        "    // Defining random points\n",
        "    for (int i = 0; i < count; i++) {\n",
        "        h_points[i].x = (float)(((rand() % 10000)) - 5000);\n",
        "        h_points[i].y = (float)(((rand() % 10000)) - 5000);\n",
        "        h_points[i].z = (float)(((rand() % 10000)) - 5000);\n",
        "    }\n",
        "\n",
        "    // Device pointers\n",
        "    int* d_indexOfClosest;\n",
        "    float3* d_points;\n",
        "\n",
        "    // Allocating memory on the device\n",
        "    cudaMalloc(&d_indexOfClosest, sizeof(int) * count);\n",
        "    cudaMalloc(&d_points, sizeof(float3) * count);\n",
        "\n",
        "    // Copying values from the host to the device\n",
        "    cudaMemcpy(d_points, h_points, sizeof(float3) * count, cudaMemcpyHostToDevice);\n",
        "\n",
        "    int threads_per_block = 64;\n",
        "    cout << \"Running brute force nearest neighbor on the GPU...\" << endl;\n",
        "    for (int i = 1; i <= 10; i++) {\n",
        "        long start = clock();\n",
        "\n",
        "        findClosestGPU<<<(count / threads_per_block) + 1, threads_per_block>>>(d_points, d_indexOfClosest, count);\n",
        "        cudaDeviceSynchronize();\n",
        "\n",
        "        // Copying results from the device to the host\n",
        "        cudaMemcpy(h_indexOfClosest, d_indexOfClosest, sizeof(int) * count, cudaMemcpyDeviceToHost);\n",
        "\n",
        "        double duration = (clock() - start) / (double)CLOCKS_PER_SEC;\n",
        "        cout << \"Test \" << i << \" took \" << duration << \" seconds\" << endl;\n",
        "    }\n",
        "\n",
        "    // Freeing device memory\n",
        "    cudaFree(d_indexOfClosest);\n",
        "    cudaFree(d_points);\n",
        "\n",
        "    // Freeing host memory\n",
        "    delete[] h_indexOfClosest;\n",
        "    delete[] h_points;\n",
        "\n",
        "    return 0;\n",
        "}\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "I8EGO9OlMwXE",
        "outputId": "3e901191-e0f6-4c87-9fd5-4bbbb0e47705"
      },
      "execution_count": 18,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting findClosestGPU.cu\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc findClosestGPU.cu -o findClosestGPU.out\n",
        "!./findClosestGPU.out"
      ],
      "metadata": {
        "id": "zMMtSLJdHi4j",
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "outputId": "f8ad5073-0192-4ba0-a7c0-9aa3d9b9f7ea"
      },
      "execution_count": 19,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Running brute force nearest neighbor on the GPU...\n",
            "Test 1 took 0.00265 seconds\n",
            "Test 2 took 0.00242 seconds\n",
            "Test 3 took 0.002423 seconds\n",
            "Test 4 took 0.002416 seconds\n",
            "Test 5 took 0.002414 seconds\n",
            "Test 6 took 0.002412 seconds\n",
            "Test 7 took 0.002409 seconds\n",
            "Test 8 took 0.002411 seconds\n",
            "Test 9 took 0.002412 seconds\n",
            "Test 10 took 0.00241 seconds\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvprof ./findClosestGPU.out"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "e0IjRDjYn21X",
        "outputId": "bd542b25-f505-4d5a-a86b-f751c0821014"
      },
      "execution_count": 20,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "==2695== NVPROF is profiling process 2695, command: ./findClosestGPU.out\n",
            "Running brute force nearest neighbor on the GPU...\n",
            "Test 1 took 0.002646 seconds\n",
            "Test 2 took 0.002424 seconds\n",
            "Test 3 took 0.002424 seconds\n",
            "Test 4 took 0.002422 seconds\n",
            "Test 5 took 0.002418 seconds\n",
            "Test 6 took 0.002423 seconds\n",
            "Test 7 took 0.00245 seconds\n",
            "Test 8 took 0.00242 seconds\n",
            "Test 9 took 0.002416 seconds\n",
            "Test 10 took 0.002421 seconds\n",
            "==2695== Profiling application: ./findClosestGPU.out\n",
            "==2695== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   99.76%  23.825ms        10  2.3825ms  2.3797ms  2.3857ms  findClosestGPU(float3*, int*, int)\n",
            "                    0.19%  45.663us        10  4.5660us  4.4160us  5.0230us  [CUDA memcpy DtoH]\n",
            "                    0.05%  12.800us         1  12.800us  12.800us  12.800us  [CUDA memcpy HtoD]\n",
            "      API calls:   77.37%  84.653ms         2  42.326ms  7.5050us  84.645ms  cudaMalloc\n",
            "                   21.82%  23.873ms        10  2.3873ms  2.3840ms  2.3912ms  cudaDeviceSynchronize\n",
            "                    0.30%  333.10us        11  30.282us  24.192us  57.187us  cudaMemcpy\n",
            "                    0.23%  251.63us        10  25.162us  5.4530us  191.80us  cudaLaunchKernel\n",
            "                    0.13%  143.57us       114  1.2590us     143ns  55.530us  cuDeviceGetAttribute\n",
            "                    0.12%  136.77us         2  68.384us  7.8910us  128.88us  cudaFree\n",
            "                    0.01%  11.536us         1  11.536us  11.536us  11.536us  cuDeviceGetName\n",
            "                    0.01%  5.8710us         1  5.8710us  5.8710us  5.8710us  cuDeviceGetPCIBusId\n",
            "                    0.00%  5.1400us         1  5.1400us  5.1400us  5.1400us  cuDeviceTotalMem\n",
            "                    0.00%  1.6550us         3     551ns     222ns  1.1750us  cuDeviceGetCount\n",
            "                    0.00%     793ns         2     396ns     168ns     625ns  cuDeviceGet\n",
            "                    0.00%     301ns         1     301ns     301ns     301ns  cuModuleGetLoadingMode\n",
            "                    0.00%     256ns         1     256ns     256ns     256ns  cuDeviceGetUuid\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [],
      "metadata": {
        "id": "-AZ6aDSJn5vP"
      },
      "execution_count": null,
      "outputs": []
    }
  ]
}